
#include <stdio.h>
#include "cush_exchange.h"
#include "cuda_util.h"
#include "cush_alloc.h"
#include "cush_free.h"

/**
 *
 * This function exchanges the data in dbuff through a temporary exchange buffer 
 * internal to the function.
 *
 **/
void cush_exchange(cush_dbuff_t* dbuff){

    if(2 != dbuff->ndevices){
        fprintf(stderr, "Inappropriate number of devices for cush_exchange\n");
        exit(EXIT_FAILURE);
    }
    
    if(dbuff->size[0] != dbuff->size[1]){
        fprintf(stderr, "Size mismatch in cush_exchange\n");
        exit(EXIT_FAILURE);
    }
    
    //set up dbuff for exchange
    cush_dbuff_t* ebuff = (cush_dbuff_t*)malloc(sizeof(cush_dbuff_t));
    cush_dp_init(ebuff, dbuff->type, 2);
    ebuff->id[0] = dbuff->id[0];
    ebuff->id[1] = dbuff->id[1];
    ebuff->size[0] = dbuff->size[0];
    ebuff->size[1] = dbuff->size[1];
    cush_alloc(ebuff);

    // perform exchange
    cush_exchange(dbuff, ebuff);

    cush_deep_free(ebuff);
}

/**
 *
 * This function exchanges the data in dbuff through a persistent exchange buffer, ebuff.
 *
 **/
void cush_exchange(cush_dbuff_t* dbuff, cush_dbuff_t* ebuff){

    if(2 != ebuff->ndevices || 2 != dbuff->ndevices){
        fprintf(stderr, "Inappropriate number of devices for cush_exchange\n");
        exit(EXIT_FAILURE);
    }
    
    if(ebuff->size[0] != ebuff->size[1] 
            || dbuff->size[0] != dbuff->size[1] 
            || ebuff->size[0] != dbuff->size[0]){
        fprintf(stderr, "Size mismatch in cush_exchange\n");
        exit(EXIT_FAILURE);
    }
    
    int i = 0;
    #pragma omp parallel for num_threads(ebuff->ndevices)
    for(i = 0; i < 2; i++){

        int idx = (i+1)%2;

        cudaSetDevice(ebuff->id[i]);

        // perform device to device exchange
        cudaMemcpyPeer(ebuff->ptr[i], i, dbuff->ptr[idx], dbuff->id[idx], ebuff->size[i]);
        checkCUDAError("exchange copy failed", ebuff->id[i]);

        // copy from the exchange buffer into the persistant buffer
        cudaMemcpy(dbuff->ptr[i], ebuff->ptr[idx], ebuff->size[idx], cudaMemcpyDeviceToDevice);
    }
}


/* vim: set sw=4 sts=4 et foldmethod=syntax syntax=c : */
